
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <cuda.h>
#include <device_functions.h>

#include <stdio.h>

#include "cv.h"
#include "highgui.h"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/core/core.hpp"
#include "opencv/cv.hpp"
#include "opencv2/core/ocl.hpp"
//#include "opencv2/core/cuda.hpp"
//#include "opencv2/cudaimgproc.hpp"

#include <iostream>
using namespace cv;
using namespace std;

typedef unsigned char uchar;
typedef unsigned int uint32;

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex;

cudaArray_t g_arr1;

__global__ void reshape_kernel(uchar * desDevImgBuff, uint32 dstWidth,uint32 dstHeight,float coeffX,float coeffY)
{
	int indexX = blockDim.x*blockIdx.x + threadIdx.x;
	int indexY = blockDim.y*blockIdx.y + threadIdx.y;

	if ((indexX >= dstWidth) || (indexY >= dstHeight))
	{
		return;
	}

	float x = float(indexX) * coeffX;
	float y = float(indexY) * coeffY;

	desDevImgBuff[indexY*dstWidth + indexX] = tex2D(tex, x, y);
}

void setDataToTexure(uchar * hostImgBuff,uint32 srcWidth,uint32 srcHeight)
{
	cudaError ret;

	cudaChannelFormatDesc chdesc = cudaCreateChannelDesc<uchar>();//cudaCreateChannelDesc<uchar>();// cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindFloat);

	ret = cudaMallocArray(&g_arr1, &chdesc, srcWidth, srcHeight);

	if (ret != cudaSuccess)
	{
		cout << "cudaMallocArray failed" << endl;
		return;
	}

	ret = cudaMemcpyToArray(g_arr1, 0, 0, hostImgBuff, srcWidth*srcHeight, cudaMemcpyHostToDevice);
	if (ret != cudaSuccess)
	{
		cout << "cudaMemcpyToArray failed" << endl;
		return;
	}

	tex.addressMode[0] = cudaAddressModeClamp;
	tex.addressMode[1] = cudaAddressModeClamp;
	tex.normalized = false;
	tex.filterMode = cudaFilterModePoint;

	ret = cudaBindTextureToArray(&tex, g_arr1, &chdesc);

	if (ret != cudaSuccess)
	{
		cout << "cudaBindTextureToArray failed" << endl;
	}
	else
	{
		cout << "cudaBindTextureToArray SUCCESS" << endl;
	}
}

void setDataToTexure_g(uchar * hostImgBuff, uint32 srcWidth, uint32 srcHeight,uchar * devBuffSrc)
{
	cudaError ret;

	cudaChannelFormatDesc chdesc = cudaCreateChannelDesc<uchar>();//cudaCreateChannelDesc<uchar>();// cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindFloat);

	cudaMemcpy(devBuffSrc, hostImgBuff, srcWidth*srcHeight,cudaMemcpyHostToDevice);

	//tex.addressMode[0] = cudaAddressModeClamp;
	//tex.addressMode[1] = cudaAddressModeClamp;
	//tex.normalized = false;
	tex.filterMode = cudaFilterModePoint;

	size_t offset = 0;
	//ret = cudaBindTexture2D(&offset, &tex, devBuffSrc, &chdesc, srcWidth, srcHeight, srcWidth);
	ret = cudaBindTexture(&offset, &tex, devBuffSrc, &chdesc, srcWidth*srcHeight);

	if (ret != cudaSuccess)
	{
		cout << "cudaBindTexture failed" << endl;
	}
	else
	{
		cout << "cudaBindTexture SUCCESS" << endl;
	}
}

__global__ void test_kernel(uchar* devBuff)
{
	devBuff[threadIdx.x] = threadIdx.x;
}

void reshape(uchar * hostImgBuffIn, uint32 widthSrc, uint32 heightSrc, uchar*hostImgBuffOut, uint32 widthDst, uint32 heightDst,uchar * dev_buff)
{
	setDataToTexure(hostImgBuffIn, widthSrc, heightSrc);

	dim3 dimBlock(32,32);
	dim3 dimGrid((widthDst + 31) / 32, (heightDst+31)/32);

	float coeffX = float(widthSrc) / float(widthDst);
	float coeffY = float(heightSrc) / float(heightDst);

	cout << "coeffX:" << coeffX << " coeffY:" << coeffY<<endl;

	reshape_kernel << <dimGrid, dimBlock >> >(dev_buff, widthDst, heightDst, coeffX, coeffY);

	//test_kernel << <1, 8 >> >(devTempBuff);
	cudaMemcpy(hostImgBuffOut, dev_buff, widthDst*heightDst, cudaMemcpyDeviceToHost);
	//cudaMemcpy(hostImgBuffOut, devTempBuff, 8, cudaMemcpyDeviceToHost);

}


#define BLOCK_NUM_TO_ALLOC 20


int main(int argc, char *argv[])
{
	uchar * g_dev_buff = NULL;

	cudaMalloc(&g_dev_buff, 1024 * 1024 * BLOCK_NUM_TO_ALLOC);

	Mat imgIn = imread(argv[1], IMREAD_GRAYSCALE);

	if (imgIn.data == NULL)
	{
		cout << "read image failed, image name: " << argv[1] << endl;
		return -1;
	}

	//cudaMemcpy(g_dev_buff, imgIn.data, imgIn.rows*imgIn.cols,cudaMemcpyHostToDevice);

	Mat imgResult = Mat::zeros(imgIn.rows+256, imgIn.cols + 256,CV_8UC1);

	reshape(imgIn.data, imgIn.cols, imgIn.rows, imgResult.data, imgResult.cols, imgResult.rows, g_dev_buff);


	//for (int i = 0; i < 8; i++)
	//{
	//	printf(" %d", imgResult.data[i]);
	//}
	cout << endl;
	namedWindow("origin", WINDOW_AUTOSIZE);
	namedWindow("reshaped", WINDOW_AUTOSIZE);

	imshow("origin", imgIn);
	imshow("reshaped", imgResult);
	waitKey(0);
	destroyWindow("origin");
	destroyWindow("reshaped");

	cudaUnbindTexture(&tex);

	cudaFreeArray(g_arr1);

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;

    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // Allocate GPU buffers for three vectors (two input, one output)    .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // Copy input vectors from host memory to GPU buffers.
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // Copy output vector from GPU buffer to host memory.
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
    
    return cudaStatus;
}
